<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<title>Jacobi: Device.cu File Reference</title>

<link href="tabs.css" rel="stylesheet" type="text/css"/>
<link href="doxygen.css" rel="stylesheet" type="text/css" />

<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>

</head>
<body>
<div id="top"><!-- do not remove this div! -->


<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  
  
  <td style="padding-left: 0.5em;">
   <div id="projectname">Jacobi
   
   </div>
   
  </td>
  
  
  
 </tr>
 </tbody>
</table>
</div>

<!-- Generated by Doxygen 1.7.6.1 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li class="current"><a href="files.html"><span>Files</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="files.html"><span>File&#160;List</span></a></li>
      <li><a href="globals.html"><span>File&#160;Members</span></a></li>
    </ul>
  </div>
</div>
<div class="header">
  <div class="summary">
<a href="#func-members">Functions</a>  </div>
  <div class="headertitle">
<div class="title">Device.cu File Reference</div>  </div>
</div><!--header-->
<div class="contents">

<p>This contains the device kernels as well as the host wrappers for these kernels.  
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &quot;<a class="el" href="Jacobi_8h.html">Jacobi.h</a>&quot;</code><br/>
</div><div class="textblock"><div class="dynheader">
Include dependency graph for Device.cu:</div>
<div class="dyncontent">
<div class="center"><img src="Device_8cu__incl.png" border="0" usemap="#Device_8cu" alt=""/></div>
<map name="Device_8cu" id="Device_8cu">
<area shape="rect" id="node3" href="Jacobi_8h.html" title="The header containing the most relevant functions for the Jacobi solver." alt="" coords="121,83,191,112"/></map>
</div>
</div><table class="memberdecls">
<tr><td colspan="2"><h2><a name="func-members"></a>
Functions</h2></td></tr>
<tr><td class="memTemplParams" colspan="2">template&lt;&gt; </td></tr>
<tr><td class="memTemplItemLeft" align="right" valign="top">__device__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="Device_8cu.html#af0933e49e05d5e8ab8f7e81ecc5e57c8">AtomicMax</a> (float *const address, const float value)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Compute the maximum of 2 single-precision floating point values using an atomic operation.  <a href="#af0933e49e05d5e8ab8f7e81ecc5e57c8"></a><br/></td></tr>
<tr><td class="memTemplParams" colspan="2">template&lt;&gt; </td></tr>
<tr><td class="memTemplItemLeft" align="right" valign="top">__device__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="Device_8cu.html#a75955732b4cf134509b8a69e2a568c05">AtomicMax</a> (double *const address, const double value)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Compute the maximum of 2 double-precision floating point values using an atomic operation.  <a href="#a75955732b4cf134509b8a69e2a568c05"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">__global__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="Device_8cu.html#aed8b87392c8fc4c4b5dd2fcd6628c229">HaloToBufferKernel</a> (real *__restrict__ const block, const real *__restrict__ const haloLeft, const real *__restrict__ const haloRight, int2 size, int hasLeftNeighbor, int hasRightNeighbor)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">The device kernel for copying (unpacking) the values from the halo buffers to the left and right side of the data block.  <a href="#aed8b87392c8fc4c4b5dd2fcd6628c229"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">__global__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="Device_8cu.html#ad7e1d75b01200c655d39f8ca01dbb9f1">BufferToHaloKernel</a> (const real *__restrict__ const block, real *__restrict__ const leftSideEdge, real *__restrict__ const rightSideEdge, int2 size, int hasLeftNeighbor, int hasRightNeighbor)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">The device kernel for copying (packing) the values on the left and right side of the data block to separate, contiguous buffers.  <a href="#ad7e1d75b01200c655d39f8ca01dbb9f1"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">__global__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="Device_8cu.html#abde66d396480d3d7cf52227a4bf204ed">JacobiComputeKernel</a> (const real *__restrict__ const oldBlock, real *__restrict__ const newBlock, int4 bounds, int stride, real *devResidue)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">The device kernel for one Jacobi iteration.  <a href="#abde66d396480d3d7cf52227a4bf204ed"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">__global__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="Device_8cu.html#ad93d62a0dd96fb1c946ae1592226d537">CopyBlockKernel</a> (const real *__restrict__ const srcBlock, real *__restrict__ const dstBlock, int4 bounds, int stride)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">The host wrapper for copying the updated block over the old one, after a Jacobi iteration finishes.  <a href="#ad93d62a0dd96fb1c946ae1592226d537"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="Device_8cu.html#a87c5ce86d74764a06a13b73dff013870">CheckCudaCall</a> (cudaError_t command, const char *commandName, const char *fileName, int line)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">The host function for checking the result of a CUDA API call.  <a href="#a87c5ce86d74764a06a13b73dff013870"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">real&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="Device_8cu.html#a08324d1c9b1ee7193242670ba4b97d88">CallJacobiKernel</a> (real *devBlocks[2], real *devResidue, const int4 *bounds, const int2 *size)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">The host wrapper for one Jacobi iteration.  <a href="#a08324d1c9b1ee7193242670ba4b97d88"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="Device_8cu.html#aeaebe085afbffaa0d6bcd77fa16ceb97">CopyDeviceBlock</a> (real *devBlocks[2], const int4 *bounds, const int2 *size)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">The host wrapper for copying the updated block over the old one, after a Jacobi iteration finishes.  <a href="#aeaebe085afbffaa0d6bcd77fa16ceb97"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="Device_8cu.html#a3b641f8b5d521a230416b64d8b809324">CopyDevHalosToBlock</a> (real *devBlock, const real *devHaloLineLeft, const real *devHaloLineRight, const int2 *size, const int *neighbors)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">The host wrapper for copying (unpacking) the values from the halo buffers to the left and right side of the data block.  <a href="#a3b641f8b5d521a230416b64d8b809324"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="Device_8cu.html#a95a52e3c2e58d2d94403a0dce2e29686">CopyDevSideEdgesFromBlock</a> (const real *devBlock, real *devSideEdges[2], const int2 *size, const int *neighbors, cudaStream_t copyStream)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">The host wrapper for copying (packing) the values on the left and right side of the data block to separate, contiguous buffers.  <a href="#a95a52e3c2e58d2d94403a0dce2e29686"></a><br/></td></tr>
</table>
<hr/><a name="details" id="details"></a><h2>Detailed Description</h2>
<div class="textblock"><p>This contains the device kernels as well as the host wrappers for these kernels. </p>
</div><hr/><h2>Function Documentation</h2>
<a class="anchor" id="af0933e49e05d5e8ab8f7e81ecc5e57c8"></a><!-- doxytag: member="Device.cu::AtomicMax" ref="af0933e49e05d5e8ab8f7e81ecc5e57c8" args="(float *const address, const float value)" -->
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;&gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ void <a class="el" href="Device_8cu.html#a75955732b4cf134509b8a69e2a568c05">AtomicMax</a> </td>
          <td>(</td>
          <td class="paramtype">float *const&#160;</td>
          <td class="paramname"><em>address</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const float&#160;</td>
          <td class="paramname"><em>value</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Compute the maximum of 2 single-precision floating point values using an atomic operation. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">address</td><td>The address of the reference value which might get updated with the maximum </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">value</td><td>The value that is compared to the reference in order to determine the maximum </td></tr>
  </table>
  </dd>
</dl>

</div>
</div>
<a class="anchor" id="a75955732b4cf134509b8a69e2a568c05"></a><!-- doxytag: member="Device.cu::AtomicMax" ref="a75955732b4cf134509b8a69e2a568c05" args="(double *const address, const double value)" -->
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;&gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ void <a class="el" href="Device_8cu.html#a75955732b4cf134509b8a69e2a568c05">AtomicMax</a> </td>
          <td>(</td>
          <td class="paramtype">double *const&#160;</td>
          <td class="paramname"><em>address</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const double&#160;</td>
          <td class="paramname"><em>value</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Compute the maximum of 2 double-precision floating point values using an atomic operation. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">address</td><td>The address of the reference value which might get updated with the maximum </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">value</td><td>The value that is compared to the reference in order to determine the maximum </td></tr>
  </table>
  </dd>
</dl>

</div>
</div>
<a class="anchor" id="ad7e1d75b01200c655d39f8ca01dbb9f1"></a><!-- doxytag: member="Device.cu::BufferToHaloKernel" ref="ad7e1d75b01200c655d39f8ca01dbb9f1" args="(const real *__restrict__ const block, real *__restrict__ const leftSideEdge, real *__restrict__ const rightSideEdge, int2 size, int hasLeftNeighbor, int hasRightNeighbor)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">__global__ void <a class="el" href="Device_8cu.html#ad7e1d75b01200c655d39f8ca01dbb9f1">BufferToHaloKernel</a> </td>
          <td>(</td>
          <td class="paramtype">const real *__restrict__ const&#160;</td>
          <td class="paramname"><em>block</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">real *__restrict__ const&#160;</td>
          <td class="paramname"><em>leftSideEdge</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">real *__restrict__ const&#160;</td>
          <td class="paramname"><em>rightSideEdge</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int2&#160;</td>
          <td class="paramname"><em>size</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>hasLeftNeighbor</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>hasRightNeighbor</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>The device kernel for copying (packing) the values on the left and right side of the data block to separate, contiguous buffers. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">block</td><td>The 2D device block containing the updated values after a Jacobi run </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">leftSideEdge</td><td>The left edge, holding the left-most updated halos </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">rightSideEdge</td><td>The right edge, holding the right-most updated halos </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">size</td><td>The 2D size of data block, excluding the edges which hold the halo values for the next iteration </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">hasLeftNeighbor</td><td>Marks if the calling MPI process has a left neighbor </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">hasRightNeighbor</td><td>Marks if the calling MPI process has a right neighbor </td></tr>
  </table>
  </dd>
</dl>

</div>
</div>
<a class="anchor" id="a08324d1c9b1ee7193242670ba4b97d88"></a><!-- doxytag: member="Device.cu::CallJacobiKernel" ref="a08324d1c9b1ee7193242670ba4b97d88" args="(real *devBlocks[2], real *devResidue, const int4 *bounds, const int2 *size)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">real <a class="el" href="Jacobi_8h.html#a08324d1c9b1ee7193242670ba4b97d88">CallJacobiKernel</a> </td>
          <td>(</td>
          <td class="paramtype">real *&#160;</td>
          <td class="paramname"><em>devBlocks</em>[2], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">real *&#160;</td>
          <td class="paramname"><em>devResidue</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const int4 *&#160;</td>
          <td class="paramname"><em>bounds</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const int2 *&#160;</td>
          <td class="paramname"><em>size</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>The host wrapper for one Jacobi iteration. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">devBlocks</td><td>The 2 blocks involved: the first is the current one, the second is the one to be updated </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">devResidue</td><td>The global residue that is to be updated through the Jacobi iteration </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">bounds</td><td>The bounds of the rectangular block region that holds only computable values </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">size</td><td>The 2D size of data blocks, excluding the edges which hold the halo values </td></tr>
  </table>
  </dd>
</dl>

<p><div class="dynheader">
Here is the call graph for this function:</div>
<div class="dyncontent">
<div class="center"><img src="Device_8cu_a08324d1c9b1ee7193242670ba4b97d88_cgraph.png" border="0" usemap="#Device_8cu_a08324d1c9b1ee7193242670ba4b97d88_cgraph" alt=""/></div>
<map name="Device_8cu_a08324d1c9b1ee7193242670ba4b97d88_cgraph" id="Device_8cu_a08324d1c9b1ee7193242670ba4b97d88_cgraph">
<area shape="rect" id="node3" href="Device_8cu.html#a87c5ce86d74764a06a13b73dff013870" title="The host function for checking the result of a CUDA API call." alt="" coords="173,5,288,35"/></map>
</div>
</p>

</div>
</div>
<a class="anchor" id="a87c5ce86d74764a06a13b73dff013870"></a><!-- doxytag: member="Device.cu::CheckCudaCall" ref="a87c5ce86d74764a06a13b73dff013870" args="(cudaError_t command, const char *commandName, const char *fileName, int line)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="Jacobi_8h.html#a87c5ce86d74764a06a13b73dff013870">CheckCudaCall</a> </td>
          <td>(</td>
          <td class="paramtype">cudaError_t&#160;</td>
          <td class="paramname"><em>command</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const char *&#160;</td>
          <td class="paramname"><em>commandName</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const char *&#160;</td>
          <td class="paramname"><em>fileName</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>line</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>The host function for checking the result of a CUDA API call. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">command</td><td>The result of the previously-issued CUDA API call </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">commandName</td><td>The name of the issued API call </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">fileName</td><td>The name of the file where the API call occurred </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">line</td><td>The line in the file where the API call occurred </td></tr>
  </table>
  </dd>
</dl>

</div>
</div>
<a class="anchor" id="ad93d62a0dd96fb1c946ae1592226d537"></a><!-- doxytag: member="Device.cu::CopyBlockKernel" ref="ad93d62a0dd96fb1c946ae1592226d537" args="(const real *__restrict__ const srcBlock, real *__restrict__ const dstBlock, int4 bounds, int stride)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">__global__ void <a class="el" href="Device_8cu.html#ad93d62a0dd96fb1c946ae1592226d537">CopyBlockKernel</a> </td>
          <td>(</td>
          <td class="paramtype">const real *__restrict__ const&#160;</td>
          <td class="paramname"><em>srcBlock</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">real *__restrict__ const&#160;</td>
          <td class="paramname"><em>dstBlock</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int4&#160;</td>
          <td class="paramname"><em>bounds</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>stride</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>The host wrapper for copying the updated block over the old one, after a Jacobi iteration finishes. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">srcBlock</td><td>The source block, from which data will be read </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">dstBlock</td><td>The destination block, where data will be written </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">bounds</td><td>The bounds of the rectangular updated region (holding only computable values) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">size</td><td>The stride of the data blocks, excluding the halo values </td></tr>
  </table>
  </dd>
</dl>

</div>
</div>
<a class="anchor" id="a3b641f8b5d521a230416b64d8b809324"></a><!-- doxytag: member="Device.cu::CopyDevHalosToBlock" ref="a3b641f8b5d521a230416b64d8b809324" args="(real *devBlock, const real *devHaloLineLeft, const real *devHaloLineRight, const int2 *size, const int *neighbors)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="Jacobi_8h.html#a3b641f8b5d521a230416b64d8b809324">CopyDevHalosToBlock</a> </td>
          <td>(</td>
          <td class="paramtype">real *&#160;</td>
          <td class="paramname"><em>devBlock</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const real *&#160;</td>
          <td class="paramname"><em>devHaloLineLeft</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const real *&#160;</td>
          <td class="paramname"><em>devHaloLineRight</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const int2 *&#160;</td>
          <td class="paramname"><em>size</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const int *&#160;</td>
          <td class="paramname"><em>neighbors</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>The host wrapper for copying (unpacking) the values from the halo buffers to the left and right side of the data block. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[out]</td><td class="paramname">devBlock</td><td>The 2D device block that will contain the halo values after unpacking </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">devHaloLineLeft</td><td>The halo buffer for the left side of the data block </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">devHaloLineRight</td><td>The halo buffer for the right side of the data block </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">size</td><td>The 2D size of data block, excluding the edges which hold the halo values </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">neighbors</td><td>The ranks of the neighboring MPI processes </td></tr>
  </table>
  </dd>
</dl>

<p><div class="dynheader">
Here is the call graph for this function:</div>
<div class="dyncontent">
<div class="center"><img src="Device_8cu_a3b641f8b5d521a230416b64d8b809324_cgraph.png" border="0" usemap="#Device_8cu_a3b641f8b5d521a230416b64d8b809324_cgraph" alt=""/></div>
<map name="Device_8cu_a3b641f8b5d521a230416b64d8b809324_cgraph" id="Device_8cu_a3b641f8b5d521a230416b64d8b809324_cgraph">
<area shape="rect" id="node3" href="Device_8cu.html#a87c5ce86d74764a06a13b73dff013870" title="The host function for checking the result of a CUDA API call." alt="" coords="213,5,328,35"/></map>
</div>
</p>

</div>
</div>
<a class="anchor" id="aeaebe085afbffaa0d6bcd77fa16ceb97"></a><!-- doxytag: member="Device.cu::CopyDeviceBlock" ref="aeaebe085afbffaa0d6bcd77fa16ceb97" args="(real *devBlocks[2], const int4 *bounds, const int2 *size)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="Jacobi_8h.html#aeaebe085afbffaa0d6bcd77fa16ceb97">CopyDeviceBlock</a> </td>
          <td>(</td>
          <td class="paramtype">real *&#160;</td>
          <td class="paramname"><em>devBlocks</em>[2], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const int4 *&#160;</td>
          <td class="paramname"><em>bounds</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const int2 *&#160;</td>
          <td class="paramname"><em>size</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>The host wrapper for copying the updated block over the old one, after a Jacobi iteration finishes. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">devBlocks</td><td>The 2 blocks involved: the first is the old one, the second is the updated one </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">bounds</td><td>The bounds of the rectangular updated region (holding only computable values) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">size</td><td>The 2D size of data blocks, excluding the edges which hold the halo values </td></tr>
  </table>
  </dd>
</dl>

<p><div class="dynheader">
Here is the call graph for this function:</div>
<div class="dyncontent">
<div class="center"><img src="Device_8cu_aeaebe085afbffaa0d6bcd77fa16ceb97_cgraph.png" border="0" usemap="#Device_8cu_aeaebe085afbffaa0d6bcd77fa16ceb97_cgraph" alt=""/></div>
<map name="Device_8cu_aeaebe085afbffaa0d6bcd77fa16ceb97_cgraph" id="Device_8cu_aeaebe085afbffaa0d6bcd77fa16ceb97_cgraph">
<area shape="rect" id="node3" href="Device_8cu.html#a87c5ce86d74764a06a13b73dff013870" title="The host function for checking the result of a CUDA API call." alt="" coords="181,5,296,35"/></map>
</div>
</p>

</div>
</div>
<a class="anchor" id="a95a52e3c2e58d2d94403a0dce2e29686"></a><!-- doxytag: member="Device.cu::CopyDevSideEdgesFromBlock" ref="a95a52e3c2e58d2d94403a0dce2e29686" args="(const real *devBlock, real *devSideEdges[2], const int2 *size, const int *neighbors, cudaStream_t copyStream)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="Jacobi_8h.html#a95a52e3c2e58d2d94403a0dce2e29686">CopyDevSideEdgesFromBlock</a> </td>
          <td>(</td>
          <td class="paramtype">const real *&#160;</td>
          <td class="paramname"><em>devBlock</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">real *&#160;</td>
          <td class="paramname"><em>devSideEdges</em>[2], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const int2 *&#160;</td>
          <td class="paramname"><em>size</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const int *&#160;</td>
          <td class="paramname"><em>neighbors</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">cudaStream_t&#160;</td>
          <td class="paramname"><em>copyStream</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>The host wrapper for copying (packing) the values on the left and right side of the data block to separate, contiguous buffers. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">devBlock</td><td>The 2D device block containing the updated values after a Jacobi run </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">devSideEdges</td><td>The buffers where the edge values will be packed in </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">size</td><td>The 2D size of data block, excluding the edges which hold the halo values for the next iteration </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">neighbors</td><td>The ranks of the neighboring MPI processes </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">copyStream</td><td>The stream on which this kernel will be executed </td></tr>
  </table>
  </dd>
</dl>

<p><div class="dynheader">
Here is the call graph for this function:</div>
<div class="dyncontent">
<div class="center"><img src="Device_8cu_a95a52e3c2e58d2d94403a0dce2e29686_cgraph.png" border="0" usemap="#Device_8cu_a95a52e3c2e58d2d94403a0dce2e29686_cgraph" alt=""/></div>
<map name="Device_8cu_a95a52e3c2e58d2d94403a0dce2e29686_cgraph" id="Device_8cu_a95a52e3c2e58d2d94403a0dce2e29686_cgraph">
<area shape="rect" id="node3" href="Device_8cu.html#a87c5ce86d74764a06a13b73dff013870" title="The host function for checking the result of a CUDA API call." alt="" coords="256,5,371,35"/></map>
</div>
</p>

</div>
</div>
<a class="anchor" id="aed8b87392c8fc4c4b5dd2fcd6628c229"></a><!-- doxytag: member="Device.cu::HaloToBufferKernel" ref="aed8b87392c8fc4c4b5dd2fcd6628c229" args="(real *__restrict__ const block, const real *__restrict__ const haloLeft, const real *__restrict__ const haloRight, int2 size, int hasLeftNeighbor, int hasRightNeighbor)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">__global__ void <a class="el" href="Device_8cu.html#aed8b87392c8fc4c4b5dd2fcd6628c229">HaloToBufferKernel</a> </td>
          <td>(</td>
          <td class="paramtype">real *__restrict__ const&#160;</td>
          <td class="paramname"><em>block</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const real *__restrict__ const&#160;</td>
          <td class="paramname"><em>haloLeft</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const real *__restrict__ const&#160;</td>
          <td class="paramname"><em>haloRight</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int2&#160;</td>
          <td class="paramname"><em>size</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>hasLeftNeighbor</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>hasRightNeighbor</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>The device kernel for copying (unpacking) the values from the halo buffers to the left and right side of the data block. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[out]</td><td class="paramname">block</td><td>The 2D device block that will contain the halo values after unpacking </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">haloLeft</td><td>The halo buffer for the left side of the data block </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">haloRight</td><td>The halo buffer for the right side of the data block </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">size</td><td>The 2D size of data block, excluding the edges which hold the halo values </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">hasLeftNeighbor</td><td>Marks if the calling MPI process has a left neighbor </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">hasRightNeighbor</td><td>Marks if the calling MPI process has a right neighbor </td></tr>
  </table>
  </dd>
</dl>

</div>
</div>
<a class="anchor" id="abde66d396480d3d7cf52227a4bf204ed"></a><!-- doxytag: member="Device.cu::JacobiComputeKernel" ref="abde66d396480d3d7cf52227a4bf204ed" args="(const real *__restrict__ const oldBlock, real *__restrict__ const newBlock, int4 bounds, int stride, real *devResidue)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">__global__ void <a class="el" href="Device_8cu.html#abde66d396480d3d7cf52227a4bf204ed">JacobiComputeKernel</a> </td>
          <td>(</td>
          <td class="paramtype">const real *__restrict__ const&#160;</td>
          <td class="paramname"><em>oldBlock</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">real *__restrict__ const&#160;</td>
          <td class="paramname"><em>newBlock</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int4&#160;</td>
          <td class="paramname"><em>bounds</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>stride</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">real *&#160;</td>
          <td class="paramname"><em>devResidue</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>The device kernel for one Jacobi iteration. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">oldBlock</td><td>The current (old) block providing the values at the start of the iteration </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">newBlock</td><td>The new block that will hold the updated values after the iteration finishes </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">bounds</td><td>The bounds of the rectangular block region that holds only computable values </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">stride</td><td>The stride of the data blocks, excluding the edges holding the halo values </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">devResidue</td><td>The global residue that is to be updated through the iteration </td></tr>
  </table>
  </dd>
</dl>

</div>
</div>
</div><!-- contents -->
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Defines</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>



<hr class="footer"/><address class="footer"><small>
Generated on Fri Mar 1 2013 15:52:54 for Jacobi by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.7.6.1
</small></address>

</body>
</html>
